Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

MV-532: [MV-MIOpen] Benchmark & Port AvgPool #48

Closed
wants to merge 30 commits into from

Conversation

hieule88
Copy link
Collaborator

@hieule88 hieule88 commented Aug 19, 2024

  • Added AvgPool 2D 3D forward and backward with new solvers and bfp16 supported.

  • Added driver test and gtest for AvgPool.

  • New API is guarded by MIOPEN_BETA_API macro.

  • Compared to ROCm pytorch: Link

  • Average over all cases:

  • AvgPool 2D

Type Forward Backward
float16 1.70 1.88
float32 1.58 2.46
bfloat16 1.77 3.34
  • AvgPool 3D
Type Forward Backward
float16 1.53 1.31
float32 1.35 Not Improved
bfloat16 1.57 1.72

@hieule88 hieule88 self-assigned this Aug 19, 2024
Copy link
Collaborator

@littlecutebird littlecutebird left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Partial review

include/miopen/miopen.h Outdated Show resolved Hide resolved
include/miopen/miopen.h Outdated Show resolved Hide resolved
src/kernels/MIOpenAvgPool.cpp Outdated Show resolved Hide resolved
src/kernels/tensor_view.hpp Outdated Show resolved Hide resolved
test/gtest/avgpool.hpp Show resolved Hide resolved
Copy link
Collaborator

@cognaiger9 cognaiger9 left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As MIOpen already has API for pooling group, you should integrate these functions into existing APIs.

@hieule88
Copy link
Collaborator Author

hieule88 commented Oct 1, 2024

As MIOpen already has API for pooling group, you should integrate these functions into existing APIs.

This was discussed with the Team leader before. We still need to port Avgpool op separately. Please continue reviewing

docs/reference/index.rst Outdated Show resolved Hide resolved
Copy link
Collaborator

@cognaiger9 cognaiger9 left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

partial review

driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/mloAvgPoolHost.hpp Outdated Show resolved Hide resolved
driver/mloAvgPoolHost.hpp Outdated Show resolved Hide resolved
driver/mloAvgPoolHost.hpp Outdated Show resolved Hide resolved
driver/mloAvgPoolHost.hpp Outdated Show resolved Hide resolved
Copy link
Collaborator

@littlecutebird littlecutebird left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just partial review until I tried to run your driver and realized that benchmark result seems to not be correct. May I ask you to check benchmark result before I continue to review? Thank you. And please run make analyze after that because I also tried to run it and the PR produced some errors, please fix those errors too.

include/miopen/miopen.h Outdated Show resolved Hide resolved
src/kernels/tensor_view.hpp Outdated Show resolved Hide resolved
src/include/miopen/avgpool.hpp Outdated Show resolved Hide resolved
src/include/miopen/tensor_view_utils.hpp Outdated Show resolved Hide resolved
@hieule88
Copy link
Collaborator Author

hieule88 commented Oct 9, 2024

Please continue reviewing @littlecutebird @cognaiger9

src/avgpool/problem_description.cpp Outdated Show resolved Hide resolved
src/avgpool/problem_description.cpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
Copy link
Collaborator

@littlecutebird littlecutebird left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Partial review (without driver and gtest)

src/avgpool_api.cpp Outdated Show resolved Hide resolved
src/include/miopen/avgpool/problem_description.hpp Outdated Show resolved Hide resolved
src/include/miopen/avgpool/problem_description.hpp Outdated Show resolved Hide resolved
src/include/miopen/avgpool/problem_description.hpp Outdated Show resolved Hide resolved
src/include/miopen/avgpool/invoke_params.hpp Outdated Show resolved Hide resolved
src/avgpool/problem_description.cpp Outdated Show resolved Hide resolved
src/solver/avgpool/backward_avgpool_2d.cpp Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
driver/avgpool_driver.hpp Outdated Show resolved Hide resolved
Copy link
Collaborator

@littlecutebird littlecutebird left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Other reviews

driver/avgpool_driver.hpp Show resolved Hide resolved
Comment on lines +70 to +83
std::vector<size_t> ComputeStrides(std::vector<size_t> inputDim) const
{
if(!is_contiguous)
std::swap(inputDim.front(), inputDim.back());
std::vector<size_t> strides(inputDim.size());
strides.back() = 1;
for(int i = inputDim.size() - 2; i >= 0; --i)
strides[i] = strides[i + 1] * inputDim[i + 1];
if(!is_contiguous)
std::swap(strides.front(), strides.back());
return strides;
}
};

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@long10024070 maybe we should put this function into an util file in gtest instead of duplicating it in every gtest files

test/gtest/avgpool.hpp Outdated Show resolved Hide resolved
test/gtest/avgpool.hpp Outdated Show resolved Hide resolved
test/gtest/avgpool.hpp Outdated Show resolved Hide resolved
test/gtest/avgpool.hpp Outdated Show resolved Hide resolved
test/gtest/avgpool.hpp Outdated Show resolved Hide resolved
@cognaiger9
Copy link
Collaborator

LGTM

@littlecutebird
Copy link
Collaborator

OK

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants